//#define fast_norm3( poot ) (poot).w = native_rsqrt(mad((poot).x, (poot).x, mad((poot).y, (poot).y, (poot).z*(poot).z))); (poot).x *= (poot).w; (poot).y *= (poot).w; (poot).z *= (poot).w; (poot).w = 0.0f;
//#define fast_norm3_t( temp, poot ) (temp) = native_rsqrt(mad((poot).x, (poot).x, mad((poot).y, (poot).y, (poot).z*(poot).z))); (poot).x *= (temp); (poot).y *= (temp); (poot).z *= (temp);

#define fast_norm3( v ) (v).w = native_rsqrt((v).x*(v).x + (v).y*(v).y + (v).z+(v).z); (v).x *= (v).w; (v).y *= (v).w; (v).z *= (v).w; (poot).w = 0.0f;
#define fast_norm3_t( temp, v ) (temp) = native_rsqrt((v).x*(v).x + (v).y*(v).y + (v).z*(v).z ); (v).x *= (temp); (v).y *= (temp); (v).z *= (temp);

#define fast_dot3( v1, v2 ) ((v1).x*(v2).x + (v1).y*(v2).y + (v1).z*(v2).z)
#define fast_sub3( r, v1, v2 ) (r).x = (v1).x - (v2).x; (r).y = (v1).y - (v2).y; (r).z = (v1).z - (v2).z;
#define fast_add3( r, v1, v2 ) (r).x = (v1).x + (v2).x; (r).y = (v1).y + (v2).y; (r).z = (v1).z + (v2).z;
#define fast_min3( r, v1, v2 ) (r).x = fmin((v1).x, (v2).x); (r).y = fmin((v1).y, (v2).y); (r).z = fmin((v1).z, (v2).z);
#define fast_max3( r, v1, v2 ) (r).x = fmax((v1).x, (v2).x); (r).y = fmax((v1).y, (v2).y); (r).z = fmax((v1).z, (v2).z);

typedef struct {
	float x;
	float y;
	float z;
} float3h;

int boxIntersect( float3h min1, float3h max1, float3h min2, float3h max2 ) {
	if ( min1.x > max2.x ) return 0;
	if ( min1.y > max2.y ) return 0;
	if ( min1.z > max2.z ) return 0;
	if ( max1.x < min2.x ) return 0;
	if ( max1.y < min2.y ) return 0;
	if ( max1.z < min2.z ) return 0;
	return 1;
}

float3h getLightPosition(float i, float animation_time) {
	i /= 16.0f;
	float3h light;
	light.x = 0.0f + (5.9f) * native_sin( -(animation_time + i)  * 10.0f);
	light.y = 0.0f;
	light.z = -10.0f + (5.9f) * native_cos( -(animation_time + i) * 10.0f );
	
	return light;
		
}

float3h project(float3h v, float Q1, float Q2, float C, float D) {
	float3h result;
	result.x = -Q1 * v.x / v.z;
	result.y = -Q2 * v.y / v.z;
	result.z = -C - D / v.z;
	return result;
}

float3h getLightColor(float i) {
	//i /= 2.0f;
	float3h color = { 0.5f * native_sin((float)i) + 0.5f, 0.5f * native_cos((float)i) + 0.5f, 0.5f * native_sin(2.0f * (float)i) + 0.5f};
	return color;
	//float3h color = { 1.0f, 1.0f, 1.0f };
	//return color;
}


__kernel void simpleGenLightBuffer(	__read_only	 image2d_t	normSpecImage,
									__read_only	 image2d_t	depthImage,
									__write_only image2d_t	diffuseImage, 
									__write_only image2d_t	specularImage,
									const		float		specularPowerScale,
									const		float		zNear,
									const		float		Q1, // 2 * zNear / ( right - left )
									const		float		Q2, // 2 * zNear / ( top - bottom )
									const		float		C,  // - (zFar + zNear) / (zFar - zNear)
									const		float		D,  // - (2 zFar zNear) / (zFar - zNear)
									const		sampler_t	nearestSampler,
									const		int2		imageSize,
									const		float		animation_time,
									__read_only image2d_t	perBlockImage1, 
									__read_only image2d_t	perBlockImage2)
{


	const float lightPower = 30.0f;


	float2 coords;
	coords.x = get_global_id(0);
	coords.y = get_global_id(1);

	if (coords.x >= imageSize.x || coords.y >= imageSize.y) {
		return;
	}
	
	int2 local_size = { get_local_size(0), get_local_size(1) };
	int2 local_id = { get_local_id(0), get_local_id(1) };

	/* per block attributes */
	float block_cone_r2;
	float3h block_min, block_max;
	float3h block_normal;
	/* end per block attributes */

	float debug1 = 0.0f;
	float debug2 = 0.0f;
	
	float temp;

	float inv_zNear = native_recip(zNear);

	__local uint block_culled[512];

	float2 inv_image_size = { native_divide(1.0f , (float)imageSize.x), native_divide(1.0f , (float)imageSize.y)  };

#if 1

	float2 block_coords = { get_group_id(0) , get_group_id(1) };

	float4 temp_lookup1;
	temp_lookup1 = read_imagef(perBlockImage1, nearestSampler, block_coords);
	block_normal.x = temp_lookup1.x;
	block_normal.y = temp_lookup1.y;
	block_normal.z = temp_lookup1.z;
	block_cone_r2  = temp_lookup1.w;
		
	float4 temp_lookup2 = read_imagef(perBlockImage2, nearestSampler, block_coords);
	block_min.z = temp_lookup2.x;
	block_max.z = temp_lookup2.y;


	float min_diff_x = local_id.x;
	float min_diff_y = local_id.y;
	float max_diff_x = local_size.x - 1.0f - local_id.x;
	float max_diff_y = local_size.y - 1.0f - local_id.y;

	block_min.x  = mad(2.0f, (coords.x - min_diff_x) * inv_image_size.x, -1.0f);
	block_min.y  = mad(2.0f, (coords.y - min_diff_y) * inv_image_size.y, -1.0f);
	block_max.x  = mad(2.0f, (coords.x + max_diff_x) * inv_image_size.x, -1.0f);
	block_max.y  = mad(2.0f, (coords.y + max_diff_y) * inv_image_size.y, -1.0f);
	
#endif


	//int num_lights = mul24( local_size.x , local_size.y );
	int num_lights = 32;

	// for some terrible reason using .x results in low precision!
	float depth = 2.0f * read_imagef(depthImage, nearestSampler, coords).x - 1.0f;

	// compute eye-sace position
	float3h worldPos;
	worldPos.x  = (2.0f * coords.x * inv_image_size.x - 1.0f) * ( depth / -Q1 );
	worldPos.y  = (2.0f * coords.y * inv_image_size.y - 1.0f) * ( depth / -Q2 );
	worldPos.z  = -D / ( depth + C);

	const float boxSize	= 6.5f;//native_sqrt(lightPower / threshold);
	
	uint linear_threadID = local_id.x + local_size.y * local_id.y;

#if 1
	uint num_threads = local_size.y * local_size.x;
	uint iterations = ( num_lights + num_threads - 1 ) / num_threads;
	uint i;
	for (i=0; i<iterations; i++) {
		uint lightIndex = linear_threadID + (num_threads * i);
		if ( lightIndex < num_lights ) {
			float3h light = getLightPosition(linear_threadID, animation_time);
			/* box test */
			
			/* project a bounding box of the light into clip space */
			if ( light.z - boxSize > -zNear ) {
				// if the backside of light's bounding box
				// is in front of the zNear plane
				// then the light should be clipped
				continue;
			}
			
			
			float3h lightMin, lightMax;
			float div1 = 1.0f / ( light.z + boxSize );
			float div2 = 1.0f / ( light.z - boxSize );
			
			if ( light.z + boxSize > -zNear ) {
				// solves the problem of when front plane of bounding box
				// is behind zNear
				div1 = -zNear;
			}
			lightMin.x = -Q1 * ( light.x - boxSize );
			lightMin.y = -Q2 * ( light.y - boxSize );
			lightMax.x = -Q1 * ( light.x + boxSize );
			lightMax.y = -Q2 * ( light.y + boxSize );
			lightMin.z = -C - D * div1;
			lightMax.z = -C - D * div2;
			if ( lightMin.x < 0.0f ) {
				lightMin.x *= div1;
				lightMax.x *= div2;
			}
			else {
				lightMin.x *= div2;
				lightMax.x *= div1;
			}
			if ( lightMin.y < 0.0f ) {
				lightMin.y *= div1;
				lightMax.y *= div2;
			}
			else {
				lightMin.y *= div2;
				lightMax.y *= div1;
			}
						
			int isCulled = !boxIntersect( block_min, block_max, lightMin, lightMax );
																											
			if (!isCulled && block_cone_r2 <= 1.0f) {
				// CONE TEST	
				float3h lightDirection;
				fast_sub3(lightDirection, light, worldPos );
				fast_norm3_t(temp, lightDirection);
				float avg_dot		= -fast_dot3(lightDirection, block_normal);
				float avg_dot_sqr	= avg_dot * avg_dot;
				float avg_dot_cub	= avg_dot_sqr * avg_dot;
				isCulled = (avg_dot > 0.0f) && ((avg_dot_sqr - avg_dot_cub) > (0.5f * block_cone_r2));
			}
			block_culled[lightIndex] = isCulled;
		}
	}
	barrier(CLK_LOCAL_MEM_FENCE);
#endif

	/* 
		Compute per pixel values that are not per light
	*/

	// lookup normal / specular
	float4 normalSpec		= read_imagef(normSpecImage, nearestSampler, coords);
	float specularPower		= normalSpec.w;
	
	// compute normal
	float3h normal			= { normalSpec.x - 0.5f, normalSpec.y - 0.5f, normalSpec.z - 0.5f};
	fast_norm3_t(temp, normal);


	// compute normalized eye position
	float3h eyeNormalized = worldPos;
	fast_norm3_t(temp, eyeNormalized);


	float specFactor = specularPower * specularPowerScale;

	float3h diffuse	= { 0.0f, 0.0f, 0.0f };
	float3h spec	= { 0.0f, 0.0f, 0.0f };


	const float box_size_sq = boxSize * boxSize;

	for (i=0; i<num_lights; i++) {
	
		if ( !block_culled[i] ) {
			
			float3h light		= getLightPosition(i, animation_time);
			float3h lightColor	= getLightColor(i);
			
			float3h lightDirection;
			fast_sub3(lightDirection, light, worldPos );
			
			float d2	= fast_dot3(lightDirection, lightDirection);
			float f		= d2 / box_size_sq;
			float klamp = 1.0f - f * f;
			float attenuation = fmin(1.0f, lightPower * klamp / d2);
			
			// normalize light direction
			fast_norm3_t(temp, lightDirection);

			// compute diffuse component
			float kd = ( attenuation * fmax( 0.0f, fast_dot3(normal, lightDirection) ) );
			diffuse.x				+= lightColor.x * kd;
			diffuse.y				+= lightColor.y * kd;
			diffuse.z				+= lightColor.z * kd;

			debug1 += 1.0f;

			// compute specular
			float3h halfV;
			fast_sub3( halfV, lightDirection, eyeNormalized);	
			fast_norm3_t(temp, halfV );
			float halfVdotN			= fmax(0.0f, fast_dot3(halfV, normal));
			float ks				= attenuation * native_powr(halfVdotN, specFactor );
			spec.x					+= lightColor.x * ks;
			spec.y					+= lightColor.y * ks;
			spec.z					+= lightColor.z * ks;
						
		}

	}
	
	debug2 = worldPos.z;
		
	int2 coordsi = { (int)coords.x, (int)coords.y };
	
	float4 out1 = { diffuse.x, diffuse.y, diffuse.y, debug1 };
	float4 out2 = { spec.x, spec.y, spec.z, debug2 };

	// write outputs
	write_imagef( diffuseImage,  coordsi, out1 );
   	write_imagef( specularImage, coordsi, out2 );

}